
__kernel void cl_memset_int(__global int * ar, int num){
	size_t stride = get_global_size(0);
	size_t offset = get_global_id(0);

	for(size_t i=offset; i<num; i+= stride)
		ar[i] = 0;
}

__kernel void count_hash_num(__global int * dim, int dNum, __global int *num){
        size_t stride = get_global_size(0); 
        size_t offset = get_global_id(0); 

        for(size_t i=offset;i<dNum;i+=stride){
                int tmp = dim[i];
                int hKey = tmp% (1024*1024*16);
                atomic_add(&(num[hKey]),1);
        }

}
__kernel void build_hash_table(__global int *dim,  int dNum, __global int *psum, __global int * bucket){

        size_t stride = get_global_size(0); 
        size_t offset = get_global_id(0); 

        for(size_t i=offset;i<dNum;i+=stride){
                int joinKey = dim[i];
                int hKey = joinKey % (1024*1024*16);
                int offset = atomic_add(&psum[hKey],1) * 2;
                bucket[offset]  = joinKey;
                offset += 1;
                int dimId = i+1;
                bucket[offset] =  dimId;
        }
}
__kernel void count_join_result(__global int * num, __global int * psum, __global int * bucket, __global int * fact, int fNum, __global int * count, __global int*factFilter){
        size_t stride = get_global_size(0); 
        size_t offset = get_global_id(0); 
        int lcount = 0;

        for(size_t i=offset;i<fNum;i+=stride){
                int fkey = fact[i];
                int hkey = fkey % (1024 * 1024*16);
                int keyNum = num[hkey];
                int pSum = psum[hkey];
		int fvalue = 0;

                for(int j=0;j<keyNum;j++){
                        int dimKey = bucket[2*j + 2*pSum];
                        int dimId = bucket[2*j + 2*pSum + 1];
                        if( dimKey == fkey){
                                lcount ++;
				fvalue = dimId;
                                break;
                        }
                }
                factFilter[i] = fvalue;
        }
        count[offset] = lcount;
}
__kernel void joinFact(__global int *resPsum, __global int * fact,  int attrSize, int  num, __global int * filter, __global int * result){
        size_t stride = get_global_size(0); 
        size_t offset = get_global_id(0); 
        int localOffset = resPsum[offset];

        for(size_t i=offset;i<num;i+=stride){
                if(filter[i] != 0){
                        result[localOffset] =  fact[i]; 
                        localOffset += 1;
                }
        }

}
__kernel void  joinDim(__global int *resPsum, __global char * dim, int attrSize, int num, __global int * filter, __global char * result){

        size_t stride = get_global_size(0); 
        size_t offset = get_global_id(0); 
        int localOffset = resPsum[offset];

        for(size_t i=offset;i<num;i+=stride){
                int dimId = filter[i];
                if( dimId != 0){
                        result[localOffset] =  dim[dimId-1]; 
                        localOffset += 1;
                }
        }

}

#define NUM_BANKS 16
#define LOG_NUM_BANKS 4

inline int CONFLICT_FREE_OFFSET(int index)
{
        return ((index) >> LOG_NUM_BANKS);
}

inline void loadSharedChunkFromMem(__local int *s_data,
                                                                           __global int *g_idata,
                                                                           int n, int baseIndex,
                                                                           int* ai, int* bi,
                                                                           int* mem_ai, int* mem_bi,
                                                                           int* bankOffsetA, int* bankOffsetB, int isNP2)
{
        size_t thid = get_local_id(0);
        *mem_ai = baseIndex + thid;
        *mem_bi = *mem_ai + get_local_size(0);

        *ai = thid;
        *bi = thid + get_local_size(0);

        // compute spacing to avoid bank conflicts
        *bankOffsetA = CONFLICT_FREE_OFFSET(*ai);
        *bankOffsetB = CONFLICT_FREE_OFFSET(*bi);

        s_data[*ai + *bankOffsetA] = g_idata[*mem_ai];

        if (isNP2)
        {
                s_data[*bi + *bankOffsetB] = (*bi < n) ? g_idata[*mem_bi] : 0;
        }
        else
        {
                s_data[*bi + *bankOffsetB] = g_idata[*mem_bi];
        }
}

inline void storeSharedChunkToMem(__global int* g_odata,
                                      __local int* s_data,
                                      int n,
                                      int ai, int bi,
                                      int mem_ai, int mem_bi,
                                      int bankOffsetA, int bankOffsetB, int isNP2)
{
    barrier(CLK_LOCAL_MEM_FENCE);

    g_odata[mem_ai] = s_data[ai + bankOffsetA];
    if (isNP2)
    {
        if (bi < n)
            g_odata[mem_bi] = s_data[bi + bankOffsetB];
    }
    else
    {
        g_odata[mem_bi] = s_data[bi + bankOffsetB];
    }
}

inline void clearLastElement(__local int* s_data,
                                 __global int *g_blockSums,
                                 int blockIndex, int storeSum)
{
    if (get_local_id(0) == 0)
    {
        int index = (get_local_size(0) << 1) - 1;
        index += CONFLICT_FREE_OFFSET(index);

        if (storeSum)
        {
            g_blockSums[blockIndex] = s_data[index];
        }

        s_data[index] = 0;
    }
}

inline int buildSum(__local int *s_data)
{
    int thid = get_local_id(0);
    int stride = 1;

    for (size_t d = get_local_size(0); d > 0; d >>= 1)
    {
        barrier(CLK_LOCAL_MEM_FENCE);

        if (thid < d)
        {
            int i  = mul24(mul24(2, stride), thid);
            int ai = i + stride - 1;
            int bi = ai + stride;

            ai += CONFLICT_FREE_OFFSET(ai);
            bi += CONFLICT_FREE_OFFSET(bi);

            s_data[bi] += s_data[ai];
        }

        stride *= 2;
    }

    return stride;
}

void scanRootToLeaves(__local int *s_data, int stride)
{
    int thid = get_local_id(0);

    for (size_t d = 1; d <= get_local_size(0); d *= 2)
    {
        stride >>= 1;

        barrier(CLK_LOCAL_MEM_FENCE);

        if (thid < d)
        {
            int i  = mul24(mul24(2, stride), thid);
            int ai = i + stride - 1;
            int bi = ai + stride;

            ai += CONFLICT_FREE_OFFSET(ai);
            bi += CONFLICT_FREE_OFFSET(bi);

            int t  = s_data[ai];
            s_data[ai] = s_data[bi];
            s_data[bi] += t;
        }
    }
}

void prescanBlock(__local int *data, int blockIndex, __global int *blockSums, int storeSum)
{
    int stride = buildSum(data);               // build the sum in place up the tree
    clearLastElement(data, blockSums,
                               (blockIndex == 0) ? get_group_id(0) : blockIndex, storeSum);
    scanRootToLeaves(data, stride);            // traverse down tree to build the scan
}

__kernel void prescan(__global int *g_odata,
                        __global int *g_idata,
                        __global int *g_blockSums,
                        int n,
                        int blockIndex,
                        int baseIndex, int storeSum, int isNP2, __local int * s_data
                                                )
{
    int ai, bi, mem_ai, mem_bi, bankOffsetA, bankOffsetB;
    int bid = get_group_id(0);
    int bsize = get_local_size(0);

    loadSharedChunkFromMem(s_data, g_idata, n,
                                  (baseIndex == 0) ?
                                  mul24(bid, (bsize << 1)):baseIndex,
                                  &ai, &bi, &mem_ai, &mem_bi,
                                  &bankOffsetA, &bankOffsetB, isNP2);

    prescanBlock(s_data, blockIndex, g_blockSums,storeSum);

    storeSharedChunkToMem(g_odata, s_data, n,
                                 ai, bi, mem_ai, mem_bi,
                                 bankOffsetA, bankOffsetB, isNP2);
}

__kernel void uniformAdd(__global int *g_data,
                           __global int *uniforms,
                           int n,
                           int blockOffset,
                           int baseIndex)
{
    __local int uni;
    if (get_local_id(0) == 0)
        uni = uniforms[get_group_id(0) + blockOffset];

    int bid = get_group_id(0);
    int bsize = get_local_size(0);

    int address = mul24(bid, (bsize << 1)) + baseIndex + get_local_id(0);

    barrier(CLK_LOCAL_MEM_FENCE);

    g_data[address]              += uni;
    g_data[address + get_local_size(0)] += (get_local_id(0) + get_local_size(0) < n) * uni;
}



